

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <cutil_inline.h>

#include "ibm_kernel.cu"


// parameters

const int THR_PER_BL = 256;
const int BLOCKS = 256;
const int DATA_N = THR_PER_BL * BLOCKS;
const int POS_SZ = DATA_N * sizeof(Pos);
const int RNG_SZ = DATA_N  * sizeof(Z);

// main program
int main(int argc, char **argv)
	{
	Pos * h_iPos, * d_iPos;
	Z * h_rngz, * d_rngz;
	unsigned * h_coll, * d_coll;
    int i;

    // use command-line specified CUDA device, otherwise use device with highest Gflops/s
    if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") )
        cutilDeviceInit(argc, argv);
    else
        cudaSetDevice( cutGetMaxGflopsDeviceId() );

    printf("Initializing data...\n");
    printf("...allocating CPU memory.\n");
    h_iPos = (Pos *)malloc(POS_SZ);
    h_rngz = (Z *)malloc(RNG_SZ);
	h_coll = (unsigned *)malloc(DATA_N*sizeof(unsigned));

    printf("...allocating GPU memory.\n");
    cutilSafeCall(cudaMalloc((void **)&d_iPos, POS_SZ));
    cutilSafeCall(cudaMalloc((void **)&d_rngz, RNG_SZ));
    cutilSafeCall(cudaMalloc((void **)&d_coll, DATA_N*sizeof(unsigned)));

    printf("...generating input data in CPU mem.\n");
    srand(123);
    //Generating input data on CPU
    for(i = 0; i < DATA_N; i++)
		{
        h_iPos[i].x = rand() % 100;
        h_iPos[i].y = rand() % 100;
        h_rngz[i].z1 = rand();
        h_rngz[i].z2 = rand();
        h_rngz[i].z3 = rand();
        h_rngz[i].z4 = rand();
        }

	clock_t start = clock();
    cutilSafeCall(cudaMemcpy(d_iPos, h_iPos, POS_SZ, cudaMemcpyHostToDevice));
    cutilSafeCall(cudaMemcpy(d_rngz, h_rngz, RNG_SZ, cudaMemcpyHostToDevice));
    cutilSafeCall(cudaMemcpy(d_coll, h_coll, DATA_N*sizeof(unsigned), 
			cudaMemcpyHostToDevice));
    cutilSafeCall( cudaThreadSynchronize() );

	for (int s=0; s<1000; s++)
		{
        scalarProdGPU<<<BLOCKS, THR_PER_BL>>>(d_iPos, d_rngz, d_coll);
        cutilCheckMsg("scalarProdGPU() execution failed\n");
        cutilSafeCall( cudaThreadSynchronize() );
		}
	
    cutilSafeCall(cudaMemcpy(h_rngz, d_rngz, RNG_SZ, cudaMemcpyDeviceToHost));
    cutilSafeCall(cudaMemcpy(h_iPos, d_iPos, POS_SZ, cudaMemcpyDeviceToHost));
    cutilSafeCall(cudaMemcpy(h_coll, d_coll, DATA_N*sizeof(unsigned), 
			cudaMemcpyDeviceToHost) );

	double dur = (double)(clock() - start)/CLOCKS_PER_SEC;
	printf("time: %f secs.\n", dur);

	//for (int i=0; i<BLOCKS; i++)
		{
		for (int j=0; j<THR_PER_BL; j++)
			printf("%u ", h_coll[1*THR_PER_BL + j]);

		printf("\n");
		}

    printf("Shutting down...\n");
    cutilSafeCall( cudaFree(d_iPos) );
    cutilSafeCall( cudaFree(d_rngz)   );
    cutilSafeCall( cudaFree(d_coll)   );
    free(h_iPos);
    free(h_rngz);
    free(h_coll);

    cudaThreadExit();

    cutilExit(argc, argv);

	}
